#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include "lib.h"

/* Verifica erros */
void cudaCheck(cudaError_t error) {
    if(error != cudaSuccess) {
        fprintf(stderr,"ERROR: %s\n", cudaGetErrorString(error));
        exit(1);
    }
}

__global__ void filter(uchar *in, uchar *out, int H, int W) {
    int i, j, t, k;

    int x = blockDim.x * blockIdx.x + threadIdx.x;
    int y = blockDim.y * blockIdx.y + threadIdx.y;

    int idx = y * W + x;
    float Pvalue = 0;

    t = FILTERSIZE - 1;
    k = FILTERDIM / 2;

    // Restringe bordas
    if (x >= W-k|| y >= H-k|| x <= k-1 || y <= k-1) return;

    // Convolucao
    for(i = 0; i <= FILTERDIM - 1; i++){
        for(j = 0; j <= FILTERDIM - 1; j++){
           Pvalue += in[(idx - k + i) + W * (j - k)];
        }
    }
    out[idx] = (uchar)(Pvalue / FILTERSIZE);
}

/* Funcao externa do C para chamar o CUDA */
extern "C" void launchCuda(PPMImage *vecIn, PPMImage *vecOut); {
    uchar *inR, *inG, *inB, *outR, *outG, *outB;
    int size = vecIn->width * vecIn->height;

    cudaCheck(cudaMalloc((void**)&inR, size * sizeof(uchar)));
    cudaCheck(cudaMalloc((void**)&inG, size * sizeof(uchar)));
    cudaCheck(cudaMalloc((void**)&inB, size * sizeof(uchar)));
    cudaCheck(cudaMalloc((void**)&outR, size * sizeof(uchar)));
    cudaCheck(cudaMalloc((void**)&outG, size * sizeof(uchar)));
    cudaCheck(cudaMalloc((void**)&outB, size * sizeof(uchar)));

    cudaCheck(cudaMemcpy(inR, vecIn->r, size * sizeof(uchar), cudaMemcpyHostToDevice));
    cudaCheck(cudaMemcpy(inG, vecIn->g, size * sizeof(uchar), cudaMemcpyHostToDevice));
    cudaCheck(cudaMemcpy(inB, vecIn->b, size * sizeof(uchar), cudaMemcpyHostToDevice));

    dim3 gridDim(vecIn->width / TILEWIDTH + 1, vecIn->height / TILEWIDTH + 1);
    dim3 blockDim(TILEWIDTH, TILEWIDTH);

    filter<<<gridDim, blockDim>>>(inR, outR, vecIn->height, vecIn->width);
    filter<<<gridDim, blockDim>>>(inG, outG, vecIn->height, vecIn->width);
    filter<<<gridDim, blockDim>>>(inB, outB, vecIn->height, vecIn->width);

    cudaCheck(cudaMemcpy(vecOut->r, outR, size * sizeof(uchar), cudaMemcpyDeviceToHost));
    cudaCheck(cudaMemcpy(vecOut->g, outG, size * sizeof(uchar), cudaMemcpyDeviceToHost));
    cudaCheck(cudaMemcpy(vecOut->b, outB, size * sizeof(uchar), cudaMemcpyDeviceToHost));

    cudaFree(inR);
    cudaFree(inB);
    cudaFree(inG);
    cudaFree(outR);
    cudaFree(outG);
    cudaFree(outB);

}

